/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/detail/atomic.h>
#include <cuqu/detail/barrier.h>
#include <cuqu/detail/types.h>

namespace cuqu
{
    namespace device
    {    
        // allocated as device memory
        struct queue {
            uint32_t eq_nr;
            //uint32_t eq_nw;
            uint32_t nevents;
            uint32_t mask;
            //raw_queue_t *raw_q; // device ptr to host pinned mem 
            //int eq_fetched;
            CUdeviceptr raw_ptr;
            detail::atomic_t errors;
            int rc;
            detail::atomic_t cnt;
            detail::barrier_t bar_head;
            detail::barrier_t bar_tail;
#if CUQU_ENABLE_PERF_HIST 
            unsigned int *hist;       
#endif

            template <typename T>
            __device__ int try_push(detail::event<T>* global_event) {
                return queue_try_push(this, global_event);
            }

            template <typename T>
            __device__ int timed_push(detail::event<T>* global_event, int timeout_ms) {
                return queue_timed_push(this, global_event, timeout_ms);
            }

            template <typename T>
            __device__ int timed_push_nobar(detail::event<T>* global_event, int timeout_ms) {
                return queue_timed_push_nobar(this, global_event, timeout_ms);
            }

            template <typename T>
            __device__ int try_fetch(detail::event<T>* global_event) {
                return queue_try_fetch(this, global_event);
            }

            template <typename T>
            __device__ int timed_fetch(detail::event<T>* global_event, int timeout_ms) {
                return queue_timed_fetch(this, global_event, timeout_ms);
            }

        };

        //typedef struct cuda_queue cuda_queue_t;

    } // end namespace device

    //typedef detail::cuda_queue *cuda_queue_t;
    typedef device::queue *cuda_queue_t;

} // end namespace cuqu

#include <cuqu/detail/device_queue_inl.h>

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
